diff --git a/README.md b/README.md index 04189246f2..d5fb37e805 100644 --- a/README.md +++ b/README.md @@ -7,7 +7,7 @@ Key features include: * HIP is very thin and has little or no performance impact over coding directly in CUDA mode. * HIP allows coding in a single-source C++ programming language including features such as templates, C++11 lambdas, classes, namespaces, and more. * HIP allows developers to use the "best" development environment and tools on each target platform. -* The [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/README.md) tools automatically convert source from CUDA to HIP. +* The [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/amd-staging/README.md) tools automatically convert source from CUDA to HIP. * Developers can specialize for the platform (CUDA or AMD) to tune for performance or handle tricky cases. New projects can be developed directly in the portable HIP C++ language and can run on either NVIDIA or AMD platforms. Additionally, HIP provides porting tools which make it easy to port existing CUDA codes to the HIP layer, with no loss of performance as compared to the original CUDA application. HIP is not intended to be a drop-in replacement for CUDA, and developers should expect to do some manual coding and performance tuning work to complete the port. diff --git a/docs/developer_guide/logging.md b/docs/developer_guide/logging.md index b712a40d7c..32a29448a1 100644 --- a/docs/developer_guide/logging.md +++ b/docs/developer_guide/logging.md @@ -85,8 +85,7 @@ ClPrint(amd::LOG_INFO, amd::LOG_INIT, "Initializing HSA stack."); ## HIP Logging Example: -Below is an example to enable HIP logging and get logging information during -execution of hipinfo, +Below is an example to enable HIP logging and get logging information during execution of hipinfo on Linux, ```console user@user-test:~/hip/bin$ export AMD_LOG_LEVEL=4 @@ -136,22 +135,7 @@ concurrentKernels: 1 cooperativeLaunch: 0 cooperativeMultiDeviceLaunch: 0 arch.hasGlobalInt32Atomics: 1 -arch.hasGlobalFloatAtomicExch: 1 -arch.hasSharedInt32Atomics: 1 -arch.hasSharedFloatAtomicExch: 1 -arch.hasFloatAtomicAdd: 1 -arch.hasGlobalInt64Atomics: 1 -arch.hasSharedInt64Atomics: 1 -arch.hasDoubles: 1 -arch.hasWarpVote: 1 -arch.hasWarpBallot: 1 -arch.hasWarpShuffle: 1 -arch.hasFunnelShift: 0 -arch.hasThreadFenceSystem: 1 -arch.hasSyncThreadsExt: 0 -arch.hasSurfaceFuncs: 0 -arch.has3dGrid: 1 -arch.hasDynamicParallelism: 0 +... gcnArch: 1012 isIntegrated: 0 maxTexture1D: 65536 @@ -178,6 +162,54 @@ memInfo.total: 7.98 GB memInfo.free: 7.98 GB (100%) ``` +On Windows, AMD_LOG_LEVEL can be set via environment variable from advanced system setting, or from Command prompt run as administrator, as shown below as an example, which shows some debug log information calling backend runtime on Windows. +``` +C:\hip\bin>set AMD_LOG_LEVEL=4 +C:\hip\bin>hipinfo +:3:C:\constructicon\builds\gfx\two\22.40\drivers\compute\vdi\device\comgrctx.cpp:33 : 605413686305 us: 29864: [tid:0x9298] Loading COMGR library. +:4:C:\constructicon\builds\gfx\two\22.40\drivers\compute\vdi\platform\runtime.cpp:83 : 605413869411 us: 29864: [tid:0x9298] init +:3:C:\constructicon\builds\gfx\two\22.40\drivers\compute\hipamd\src\hip_context.cpp:47 : 605413869502 us: 29864: [tid:0x9298] Direct Dispatch: 0 +:3:C:\constructicon\builds\gfx\two\22.40\drivers\compute\hipamd\src\hip_device_runtime.cpp:543 : 605413870553 us: 29864: [tid:0x9298] hipGetDeviceCount: Returned hipSuccess : +:3:C:\constructicon\builds\gfx\two\22.40\drivers\compute\hipamd\src\hip_device_runtime.cpp:556 : 605413870631 us: 29864: [tid:0x9298] ←[32m hipSetDevice ( 0 ) ←[0m +:3:C:\constructicon\builds\gfx\two\22.40\drivers\compute\hipamd\src\hip_device_runtime.cpp:561 : 605413870848 us: 29864: [tid:0x9298] hipSetDevice: Returned hipSuccess : +-------------------------------------------------------------------------------- +device# 0 +:3:C:\constructicon\builds\gfx\two\22.40\drivers\compute\hipamd\src\hip_device.cpp:346 : 605413871623 us: 29864: [tid:0x9298] ←[32m hipGetDeviceProperties ( 0000008AEBEFF8C8, 0 ) ←[0m +:3:C:\constructicon\builds\gfx\two\22.40\drivers\compute\hipamd\src\hip_device.cpp:348 : 605413871695 us: 29864: [tid:0x9298] hipGetDeviceProperties: Returned hipSuccess : +Name: AMD Radeon(TM) Graphics +pciBusID: 3 +pciDeviceID: 0 +pciDomainID: 0 +multiProcessorCount: 7 +maxThreadsPerMultiProcessor: 2560 +isMultiGpuBoard: 0 +clockRate: 1600 Mhz +memoryClockRate: 1333 Mhz +memoryBusWidth: 0 +totalGlobalMem: 12.06 GB +totalConstMem: 2147483647 +sharedMemPerBlock: 64.00 KB +... +gcnArchName: gfx90c:xnack- +:3:C:\constructicon\builds\gfx\two\22.40\drivers\compute\hipamd\src\hip_device_runtime.cpp:541 : 605413924779 us: 29864: [tid:0x9298] ←[32m hipGetDeviceCount ( 0000008AEBEFF8A4 ) ←[0m +:3:C:\constructicon\builds\gfx\two\22.40\drivers\compute\hipamd\src\hip_device_runtime.cpp:543 : 605413925075 us: 29864: [tid:0x9298] hipGetDeviceCount: Returned hipSuccess : +peers: :3:C:\constructicon\builds\gfx\two\22.40\drivers\compute\hipamd\src\hip_peer.cpp:176 : 605413928643 us: 29864: [tid:0x9298] ←[32m hipDeviceCanAccessPeer ( 0000008AEBEFF890, 0, 0 ) ←[0m +:3:C:\constructicon\builds\gfx\two\22.40\drivers\compute\hipamd\src\hip_peer.cpp:177 : 605413928743 us: 29864: [tid:0x9298] hipDeviceCanAccessPeer: Returned hipSuccess : +non-peers: :3:C:\constructicon\builds\gfx\two\22.40\drivers\compute\hipamd\src\hip_peer.cpp:176 : 605413930830 us: 29864: [tid:0x9298] ←[32m hipDeviceCanAccessPeer ( 0000008AEBEFF890, 0, 0 ) ←[0m +:3:C:\constructicon\builds\gfx\two\22.40\drivers\compute\hipamd\src\hip_peer.cpp:177 : 605413930882 us: 29864: [tid:0x9298] hipDeviceCanAccessPeer: Returned hipSuccess : +device#0 +... +:4:C:\constructicon\builds\gfx\two\22.40\drivers\compute\vdi\device\pal\palmemory.cpp:430 : 605414517802 us: 29864: [tid:0x9298] Free-: 8000 bytes, VM[ 3007c8000, 3007d0000] +:3:C:\constructicon\builds\gfx\two\22.40\drivers\compute\vdi\device\devprogram.cpp:2979: 605414517893 us: 29864: [tid:0x9298] For Init/Fini: Kernel Name: __amd_rocclr_copyBufferToImage +:3:C:\constructicon\builds\gfx\two\22.40\drivers\compute\vdi\device\devprogram.cpp:2979: 605414518259 us: 29864: [tid:0x9298] For Init/Fini: Kernel Name: __amd_rocclr_copyBuffer +... +:4:C:\constructicon\builds\gfx\two\22.40\drivers\compute\vdi\device\pal\palmemory.cpp:206 : 605414523422 us: 29864: [tid:0x9298] Alloc: 100000 bytes, ptr[00000003008D0000-00000003009D0000], obj[00000003007D0000-00000003047D0000] +:4:C:\constructicon\builds\gfx\two\22.40\drivers\compute\vdi\device\pal\palmemory.cpp:206 : 605414523767 us: 29864: [tid:0x9298] Alloc: 100000 bytes, ptr[00000003009D0000-0000000300AD0000], obj[00000003007D0000-00000003047D0000] +:3:C:\constructicon\builds\gfx\two\22.40\drivers\compute\hipamd\src\hip_memory.cpp:681 : 605414524092 us: 29864: [tid:0x9298] hipMemGetInfo: Returned hipSuccess : +memInfo.total: 12.06 GB +memInfo.free: 11.93 GB (99%) +``` + ## HIP Logging Tips: - HIP logging works for both release and debug version of HIP application. diff --git a/docs/how_to_guides/debugging.md b/docs/how_to_guides/debugging.md index 556a51c106..8564586a86 100644 --- a/docs/how_to_guides/debugging.md +++ b/docs/how_to_guides/debugging.md @@ -100,7 +100,7 @@ Reading symbols from ./hipTexObjPitch... (gdb) break main Breakpoint 1 at 0x4013d1: file /home/test/hip/tests/src/texture/hipTexObjPitch.cpp, line 98. (gdb) run -Starting program: /home/test/hip/build/directed_tests/texture/hipTexObjPitch +Starting program: /home/test/hip/build/directed_tests/texture/hipTexObjPitch [Thread debugging using libthread_db enabled] Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1". @@ -112,11 +112,11 @@ Breakpoint 1, main () ``` ### Other Debugging Tools -There are also other debugging tools available online developers can google and choose the one best suits the debugging requirements. +There are also other debugging tools available online developers can google and choose the one best suits the debugging requirements. For example, Microsoft Visual Studio and Windgb tools are options on Windows. ## Debugging HIP Applications -Below is an example to show how to get useful information from the debugger while running a simple memory copy test, which caused an issue of segmentation fault. +Below is an example on Linux to show how to get useful information from the debugger while running a simple memory copy test, which caused an issue of segmentation fault. ```console test: simpleTest2 numElements=4194304 sizeElements=4194304 bytes @@ -176,11 +176,14 @@ Thread 1 "hipMemcpy_simpl" received signal SIGSEGV, Segmentation fault. ... ``` +On Windows, debugging HIP applications on IDE like Microsoft Visual Studio tools, are more informative and visible to debug codes, inspect variables, watch multiple details and examine the call stacks. + ## Useful Environment Variables -HIP provides some environment variables which allow HIP, hip-clang, or HSA driver to disable some feature or optimization. + +HIP provides some environment variables which allow HIP, hip-clang, or HSA driver on Linux to disable some feature or optimization. These are not intended for production but can be useful diagnose synchronization problems in the application (or driver). -Some of the most useful environment variables are described here. They are supported on the ROCm path. +Some of the most useful environment variables are described here. They are supported on the ROCm path on Linux and Windows as well. ### Kernel Enqueue Serialization Developers can control kernel command serialization from the host using the environment variable, @@ -221,8 +224,8 @@ if (totalDeviceNum > 2) { Developers can dump code object to analyze compiler related issues via setting environment variable, GPU_DUMP_CODE_OBJECT -### HSA related environment variables -HSA provides some environment variables help to analyze issues in driver or hardware, for example, +### HSA related environment variables on Linux +On Linux with open source, HSA provides some environment variables help to analyze issues in driver or hardware, for example, HSA_ENABLE_SDMA=0 It causes host-to-device and device-to-host copies to use compute shader blit kernels rather than the dedicated DMA copy engines. @@ -246,12 +249,12 @@ The following is the summary of the most useful environment variables in HIP. | AMD_SERIALIZE_KERNEL
Serialize kernel enqueue. | 0 | 1: Wait for completion before enqueue.
2: Wait for completion after enqueue.
3: Both. | | AMD_SERIALIZE_COPY
Serialize copies. | 0 | 1: Wait for completion before enqueue.
2: Wait for completion after enqueue.
3: Both. | | HIP_HOST_COHERENT
Coherent memory in hipHostMalloc. | 0 | 0: memory is not coherent between host and GPU.
1: memory is coherent with host. | -| AMD_DIRECT_DISPATCH
Enable direct kernel dispatch. | 1 | 0: Disable.
1: Enable. | +| AMD_DIRECT_DISPATCH
Enable direct kernel dispatch (Currently for Linux, under development on Windows). | 1 | 0: Disable.
1: Enable. | | GPU_MAX_HW_QUEUES
The maximum number of hardware queues allocated per device. | 4 | The variable controls how many independent hardware queues HIP runtime can create per process, per device. If application allocates more HIP streams than this number, then HIP runtime will reuse the same hardware queues for the new streams in round robin manner. Please note, this maximum number does not apply to either hardware queues that are created for CU masked HIP streams, or cooperative queue for HIP Cooperative Groups (there is only one single queue per device). | ## General Debugging Tips - 'gdb --args' can be used to conveniently pass the executable and arguments to gdb. -- From inside GDB, you can set environment variables "set env". Note the command does not use an '=' sign: +- From inside GDB on Linux, you can set environment variables "set env". Note the command does not use an '=' sign: ``` (gdb) set env AMD_SERIALIZE_KERNEL 3 diff --git a/docs/reference/kernel_language.md b/docs/reference/kernel_language.md index 3cd4d55007..6224449ab6 100644 --- a/docs/reference/kernel_language.md +++ b/docs/reference/kernel_language.md @@ -126,15 +126,15 @@ The `__restrict__` keyword tells the compiler that the associated memory pointer ## Built-In Variables -(coordinate_builtins)= ### Coordinate Built-Ins Built-ins determine the coordinate of the active work item in the execution grid. They are defined in amd_hip_runtime.h (rather than being implicitly defined by the compiler). In HIP, built-ins coordinate variable definitions are the same as in Cuda, for instance: threadIdx.x, blockIdx.y, gridDim.y, etc. The products gridDim.x * blockDim.x, gridDim.y * blockDim.y and gridDim.z * blockDim.z are always less than 2^32. +Coordinates builtins are implemented as structures for better performance. When used with printf, they needs to be casted to integer types explicitly. ### warpSize -The warpSize variable is of type int and contains the warp size (in threads) for the target device. Note that all current Nvidia devices return 32 for this variable, and all current AMD devices return 64. Device code should use the warpSize built-in to develop portable wave-aware code. +The warpSize variable is of type int and contains the warp size (in threads) for the target device. Note that all current Nvidia devices return 32 for this variable, and current AMD devices return 64 for gfx9 and 32 for gfx10 and above. The warpSize variable should only be used in device functions. Device code should use the warpSize built-in to develop portable wave-aware code. ## Vector Types diff --git a/docs/user_guide/faq.md b/docs/user_guide/faq.md index ab67cc745f..7c5870ac36 100644 --- a/docs/user_guide/faq.md +++ b/docs/user_guide/faq.md @@ -148,6 +148,9 @@ ROCclr (Radeon Open Compute Common Language Runtime) is a virtual device interfa ## What is HIPAMD? HIPAMD is a repository branched out from HIP, mainly the implementation for AMD GPU. +## Can I get HIP open source repository for Windows? +No, there is no HIP repository open publicly on Windows. + ## Can a HIP binary run on both AMD and Nvidia platforms? HIP is a source-portable language that can be compiled to run on either AMD or NVIDIA platform. HIP tools don't create a "fat binary" that can run on either platform, however. @@ -237,6 +240,11 @@ Once source is compiled with per-thread default stream enabled, all APIs will be Besides, per-thread default stream be enabled per translation unit, users can compile some files with feature enabled and some with feature disabled. Feature enabled translation unit will have default stream as per thread and there will not be any implicit synchronization done but other modules will have legacy default stream which will do implicit synchronization. +## Can I develop applications with HIP APIs on Windows the same on Linux? + +Yes, HIP APIs are available to use on both Linux and Windows. +Due to different working mechanisms on operating systems like Windows vs Linux, HIP APIs call corresponding lower level backend runtime libraries and kernel drivers for the OS, in order to control the executions on GPU hardware accordingly. There might be a few differences on the related backend software and driver support, which might affect usage of HIP APIs. See OS support details in HIP API document. + ## How can I know the version of HIP? HIP version definition has been updated since ROCm 4.2 release as the following: diff --git a/docs/user_guide/programming_manual.md b/docs/user_guide/programming_manual.md index 012a6a24f0..827a62af67 100644 --- a/docs/user_guide/programming_manual.md +++ b/docs/user_guide/programming_manual.md @@ -14,10 +14,9 @@ GPU can directly access the host memory over the CPU/GPU interconnect, without n There are flags parameter which can specify options how to allocate the memory, for example, hipHostMallocPortable, the memory is considered allocated by all contexts, not just the one on which the allocation is made. hipHostMallocMapped, will map the allocation into the address space for the current device, and the device pointer can be obtained with the API hipHostGetDevicePointer(). -hipHostMallocNumaUser is the flag to allow host memory allocation to follow numa policy by user. -All allocation flags are independent, and can be used in any combination without restriction, for instance, hipHostMalloc can be called with both hipHostMallocPortable and hipHostMallocMapped flags set. Both usage models described above use the same allocation flags, and the difference is in how the surrounding code uses the host memory. +hipHostMallocNumaUser is the flag to allow host memory allocation to follow Numa policy by user. Please note this flag is currently only applicable on Linux, under development on Windows. -See the hipHostMalloc API for more information. +All allocation flags are independent, and can be used in any combination without restriction, for instance, hipHostMalloc can be called with both hipHostMallocPortable and hipHostMallocMapped flags set. Both usage models described above use the same allocation flags, and the difference is in how the surrounding code uses the host memory. ### Numa-aware host memory allocation Numa policy determines how memory is allocated. @@ -25,44 +24,8 @@ Target of Numa policy is to select a CPU that is closest to each GPU. Numa distance is the measurement of how far between GPU and CPU devices. By default, each GPU selects a Numa CPU node that has the least Numa distance between them, that is, host memory will be automatically allocated closest on the memory pool of Numa node of the current GPU device. Using hipSetDevice API to a different GPU will still be able to access the host allocation, but can have longer Numa distance. +Note, Numa policy is so far implemented on Linux, and under development on Windows. -### Managed memory allocation -Managed memory, including the `__managed__` keyword, is supported in HIP combined host/device compilation. - -Managed memory, via unified memory allocation, allows data be shared and accessible to both the CPU and GPU using a single pointer. -The allocation will be managed by AMD GPU driver using the linux HMM (Heterogeneous Memory Management) mechanism, the user can call managed memory API hipMallocManaged to allocate a large chuch of HMM memory, execute kernels on device and fetch data between the host and device as needed. - -In HIP application, It is recommend to do the capability check before calling the managed memory APIs. For example: - -``` -int managed_memory = 0; -HIPCHECK(hipDeviceGetAttribute(&managed_memory, - hipDeviceAttributeManagedMemory,p_gpuDevice)); - -if (!managed_memory ) { - printf ("info: managed memory access not supported on the device %d\n Skipped\n", p_gpuDevice); -} -else { - HIPCHECK(hipSetDevice(p_gpuDevice)); - HIPCHECK(hipMallocManaged(&Hmm, N * sizeof(T))); -. . . -} -``` -Please note, the managed memory capability check may not be necessary, but if HMM is not supported, then managed malloc will fall back to using system memory and other managed memory API calls will have undefined behavior. -For more details on managed memory APIs, please refer to the documentation HIP-API.pdf, and the application at (https://github.com/ROCm-Developer-Tools/HIP/blob/rocm-4.5.x/tests/src/runtimeApi/memory/hipMallocManaged.cpp) is a sample usage. - -### HIP Stream Memory Operations - -HIP supports Stream Memory Operations to enable direct synchronization between Network Nodes and GPU. Following new APIs are added, - hipStreamWaitValue32 - hipStreamWaitValue64 - hipStreamWriteValue32 - hipStreamWriteValue64 - -Note, CPU access to the semaphore's memory requires volatile keyword to disable CPU compiler's optimizations on memory access. -For more details, please check the documentation HIP-API.pdf. - -Please note, HIP stream does not gurantee concurrency on AMD hardware for the case of multiple (at least 6) long running streams executing concurrently, using hipStreamSynchronize(nullptr) for synchronization. ### Coherency Controls ROCm defines two coherency options for host memory: @@ -107,8 +70,49 @@ A stronger system-level fence can be specified when the event is created with hi - Coherent host memory is the default and is the easiest to use since the memory is visible to the CPU at typical synchronization points. This memory allows in-kernel synchronization commands such as threadfence_system to work transparently. - HIP/ROCm also supports the ability to cache host memory in the GPU using the "Non-Coherent" host memory allocations. This can provide performance benefit, but care must be taken to use the correct synchronization. +### Managed memory allocation +Managed memory, including the `__managed__` keyword, is supported in HIP combined host/device compilation, on Linux, not on Windows (under development). + +Managed memory, via unified memory allocation, allows data be shared and accessible to both the CPU and GPU using a single pointer. +The allocation will be managed by AMD GPU driver using the linux HMM (Heterogeneous Memory Management) mechanism, the user can call managed memory API hipMallocManaged to allocate a large chuch of HMM memory, execute kernels on device and fetch data between the host and device as needed. + +In HIP application, It is recommend to do the capability check before calling the managed memory APIs. For example: + +``` +int managed_memory = 0; +HIPCHECK(hipDeviceGetAttribute(&managed_memory, + hipDeviceAttributeManagedMemory,p_gpuDevice)); + +if (!managed_memory ) { + printf ("info: managed memory access not supported on the device %d\n Skipped\n", p_gpuDevice); +} +else { + HIPCHECK(hipSetDevice(p_gpuDevice)); + HIPCHECK(hipMallocManaged(&Hmm, N * sizeof(T))); +. . . +} +``` +Please note, the managed memory capability check may not be necessary, but if HMM is not supported, then managed malloc will fall back to using system memory and other managed memory API calls will have undefined behavior. +For more details on managed memory APIs, please refer to the documentation HIP-API.pdf, and the application at (https://github.com/ROCm-Developer-Tools/HIP/blob/rocm-4.5.x/tests/src/runtimeApi/memory/hipMallocManaged.cpp) is a sample usage. + +Note, managed memory management is implemented on Linux, not supported on Windows yet. + +### HIP Stream Memory Operations + +HIP supports Stream Memory Operations to enable direct synchronization between Network Nodes and GPU. Following new APIs are added, + hipStreamWaitValue32 + hipStreamWaitValue64 + hipStreamWriteValue32 + hipStreamWriteValue64 + +Note, CPU access to the semaphore's memory requires volatile keyword to disable CPU compiler's optimizations on memory access. +For more details, please check the documentation HIP-API.pdf. + +Please note, HIP stream does not gurantee concurrency on AMD hardware for the case of multiple (at least 6) long running streams executing concurrently, using hipStreamSynchronize(nullptr) for synchronization. + ## Direct Dispatch -HIP runtime has Direct Dispatch enabled by default in ROCM 4.4. With this feature we move away from our conventional producer-consumer model where the runtime creates a worker thread(consumer) for each HIP Stream, and the host thread(producer) enqueues commands to a command queue(per stream). +HIP runtime has Direct Dispatch enabled by default in ROCM 4.4 on Linux. +With this feature we move away from our conventional producer-consumer model where the runtime creates a worker thread(consumer) for each HIP Stream, and the host thread(producer) enqueues commands to a command queue(per stream). For Direct Dispatch, HIP runtime would directly enqueue a packet to the AQL queue (user mode queue on GPU) on the Dispatch API call from the application. That has shown to reduce the latency to launch the first wave on the idle GPU and total time of tiny dispatches synchronized with the host. @@ -117,14 +121,15 @@ In addition, eliminating the threads in runtime has reduced the variance in the This feature can be disabled via setting the following environment variable, AMD_DIRECT_DISPATCH=0 +Note, Direct Dispatch is implemented on Linux. It is currently not supported on Windows. + ## HIP Runtime Compilation HIP now supports runtime compilation (hipRTC), the usage of which will provide the possibility of optimizations and performance improvement compared with other APIs via regular offline static compilation. hipRTC APIs accept HIP source files in character string format as input parameters and create handles of programs by compiling the HIP source files without spawning separate processes. -For more details on hipRTC APIs, refer to HIP-API.pdf in GitHub (https://github.com/RadeonOpenCompute/ROCm). - -The link here(https://github.com/ROCm-Developer-Tools/HIP/blob/main/tests/src/hiprtc/saxpy.cpp) shows an example how to program HIP application using runtime compilation mechanism, and detail hipRTC programming guide is also available in Github (https://github.com/ROCm-Developer-Tools/HIP/blob/main/docs/markdown/hip_rtc.md). +For more details on hipRTC APIs, refer to HIP-API.pdf in GitHub (https://docs.amd.com/category/api_documentation). +For Linux developers, the link here(https://github.com/ROCm-Developer-Tools/hip-tests/blob/develop/samples/2_Cookbook/23_cmake_hiprtc/saxpy.cpp) shows an example how to program HIP application using runtime compilation mechanism, and detail hipRTC programming guide is also available in Github (https://github.com/ROCm-Developer-Tools/HIP/blob/develop/docs/user_guide/hip_rtc.md). ## HIP Graph HIP graph is supported. For more details, refer to the HIP API Guide. @@ -143,7 +148,7 @@ The per-thread default stream is a blocking stream and will synchronize with the The per-thread default stream can be enabled via adding a compilation option, “-fgpu-default-stream=per-thread”. -And users can explicitly use "hipStreamPerThread" as per-thread default stream handle as input in API commands. There are test codes as examples in the link (https://github.com/ROCm-Developer-Tools/HIP/tree/develop/tests/catch/unit/streamperthread). +And users can explicitly use "hipStreamPerThread" as per-thread default stream handle as input in API commands. There are test codes as examples in the link (https://github.com/ROCm-Developer-Tools/hip-tests/tree/develop/catch/unit/streamperthread). ## Use of Long Double Type