backup temp changes

CMakeLists.txt

@@ -234,7 +234,7 @@ set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/lib)
 set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/lib)
 set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/bin)
 add_subdirectory(src)
-add_subdirectory(doc)
+#add_subdirectory(doc)
 add_subdirectory(test)
 add_subdirectory(tools)
 

src/include/migraphx/shape.hpp

@@ -223,11 +223,11 @@ struct shape
     static type_t parse_type(const std::string& s);
 
     const std::vector<shape>& sub_shapes() const;
+    std::size_t element_space() const;
 
     private:
     std::shared_ptr<const shape_impl> impl;
 
-    std::size_t element_space() const;
 };
 
 void migraphx_to_value(value& v, const shape& s);

src/targets/gpu/device/mul_add.cpp

@@ -21,7 +21,7 @@ namespace device {
 //    }
 //}
 
-__global__ void mul_add_kernel(void* a, void* x, void* b, void* r, int n)
+__global__ void mul_add_kernel(void* a, int an, void* x, int xn, void* b, int bn, void* r, int n)
 {
     int id = blockDim.x * blockIdx.x + threadIdx.x;
     __half2* ha = reinterpret_cast<__half2*>(a);
@@ -30,7 +30,7 @@ __global__ void mul_add_kernel(void* a, void* x, void* b, void* r, int n)
     __half2* hr = reinterpret_cast<__half2*>(r);
     if (id < n)
     {
-        hr[id] = __hadd2(__hmul2(ha[id], hx[id]), hb[id]);
+        hr[id] = __hadd2(__hmul2(ha[id % an], hx[id % xn]), hb[id % bn]);
     }
 }
 
@@ -44,13 +44,22 @@ void mul_add(hipStream_t stream,
     auto type = result.get_shape().type();
     if (type == shape::half_type)
     {
-        auto elem_num = result.get_shape().elements() / 2;
+        std::cout << "case1" << std::endl;
+        int s1e = arg1.get_shape().element_space() / 2;
+        int s2e = arg2.get_shape().element_space() / 2;
+        int s3e = arg3.get_shape().element_space() / 2;
+        int elem_num = result.get_shape().elements() / 2;
+        s1e = (s1e == 0 ? 1 : s1e);
+        s2e = (s2e == 0 ? 1 : s2e);
+        s3e = (s3e == 0 ? 1 : s3e);
+        std::cout << "re =" << elem_num << ", s1e = " << s1e << ", s2e = " << s2e << ", s3e = " << s3e << std::endl;
         int block_size = 1024;
         int block_num = (elem_num + block_size - 1) / block_size;
-        mul_add_kernel<<<block_num, block_size>>>(arg1.data(), arg2.data(), arg3.data(), result.data(), elem_num);
+        mul_add_kernel<<<block_num, block_size>>>(arg1.data(), s1e, arg2.data(), s2e, arg3.data(), s3e, result.data(), elem_num);
     }
     else
     {
+        std::cout << "case2" << std::endl;
         nary(stream, result, arg1, arg2, arg3)([](auto x, auto a, auto b)
                                                 __device__ { return a * x + b; });        
     }