Update square sample for recent HIP ease-of-use improvements

[ROCm/hip commit: b9fa704521]
Этот коммит содержится в:
Ben Sander
2017-12-02 07:44:27 -06:00
родитель cccca66ce4
Коммит 2bd4446721
3 изменённых файлов: 16 добавлений и 22 удалений
+9 -12
Просмотреть файл
@@ -1,20 +1,17 @@
HIP_PATH?= $(wildcard /opt/rocm/hip)
HIPCC=$(HIP_PATH)/bin/hipcc
all: square.hip.out
all: square.out
# Step
square.cpp:
$(HIP_PATH)/bin/hipify-perl square.cu > square.cpp
square.out: square.cpp
$(HIPCC) $(CXXFLAGS) square.cpp -o $@
# make step to make native cuda app, if on cuda machine
square.cuda.out : square.cu
nvcc square.cu -o $@
#hipify square.cu > square.cpp
# Then review & finish port in square.cpp
#
square.hip.out: square.hipref.cpp
$(HIPCC) $(CXXFLAGS) square.hipref.cpp -o $@
clean:
rm -f *.o *.out
+6 -9
Просмотреть файл
@@ -1,16 +1,13 @@
# Square.md
Simple test which shows how to use hipify to port CUDA code to HIP.
See related [blog](http://gpuopen.com/hip-to-be-squared-an-introductory-hip-tutorial) that explains the example.
See related [blog](http://gpuopen.com/hip-to-be-squared-an-introductory-hip-tutorial) that explains the example.
Now it is even simpler and requires no manual modification to the hipified source code - just hipify and compile:
1. Add hip/bin path to the PATH :
<code>export PATH=$PATH:[MYHIP]/bin</code>
2. Do <code>$ hipify square.cu > square.cpp </code>
3. Manually edit square.cpp to add hipLaunchParms lp to kernel parms:
<code>vector_square(hipLaunchParm lp, T *C_d, const T *A_d, size_t N)</code>
(see square.hipref.cpp for the correct output after running hipify and the above manual step)
4. make
2. <code>$ make </code>
Make runs these steps. This can be performed on either CUDA or AMD platform:
<code>hipify-perl square.cu > square.cpp </code> # convert cuda code to hip code
<code>hipcc square.cpp</code> # compile into executable
+1 -1
Просмотреть файл
@@ -38,7 +38,7 @@ THE SOFTWARE.
*/
template <typename T>
__global__ void
vector_square(T *C_d, const T *A_d, size_t N)
vector_square(T *C_d, T *A_d, size_t N)
{
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x ;