diff --git a/docs/markdown/hip_porting_guide.md b/docs/markdown/hip_porting_guide.md index cc964bc62a..621726ee5f 100644 --- a/docs/markdown/hip_porting_guide.md +++ b/docs/markdown/hip_porting_guide.md @@ -402,61 +402,54 @@ Code should not assume a warp size of 32 or 64. See [Warp Cross-Lane Functions] ## memcpyToSymbol -HIP support for hipMemCpyToSymbol is under-development. This feature allows a kernel +HIP support for hipMemCpyToSymbol is complete. This feature allows a kernel to define a device-side data symbol which can be accessed on the host side. The symbol -can be in __constant or device space. As a workaround, programs can pass the symbol -as an argument to the kernel, and use standard hipMemcpy routines to initialize it. +can be in __constant or device space. For example: Device Code: ``` -#include -#include +#include +#include #include -#ifdef __HIP_PLATFORM_HCC__ -__global__ void Inc(hipLaunchParm lp, float *Ad, float *Out) -#endif -#ifdef __HIP_PLATFORM_NVCC__ -__constant__ float Ad[1024]; -__global__ void Inc(hipLaunchParm lp, float *Out) -#endif +#define HIP_ASSERT(status) \ + assert(status == hipSuccess) + +#define LEN 512 +#define SIZE 2048 + +__constant__ int Value[LEN]; + +__global__ void Get(hipLaunchParm lp, int *Ad) { - int tx = hipThreadIdx_x; - Out[tx] = Ad[tx] + 1.0f; + int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + Ad[tid] = Value[tid]; } int main() { - float *A, *Ad; - float *Out, *Outd; - A = new float[1024]; - Out = new float[1024]; - - for(uint32_t i=0;i<1024;i++) + int *A, *B, *Ad; + A = new int[LEN]; + B = new int[LEN]; + for(unsigned i=0;i