revert workaround for square sample and update doc on GGL
Change-Id: I731c68ca4111e7dc2e45bef51c4cad2c23fc81f8
[ROCm/clr commit: 14f8c7cf74]
Esse commit está contido em:
@@ -243,10 +243,4 @@ $HIP/include/hip/hcc_detail/hip_runtime_api.h
|
||||
$HIP/include/hip/hcc_detail/host_defines.h
|
||||
Or pass "-DGENERIC_GRID_LAUNCH=1" to hipcc at application compilation time.
|
||||
|
||||
There are some limitation/assumptions of GGL implementation right now:
|
||||
1. GGL was only tested with Ubuntu16.04, assuming HCC and HIP only link against libstdc++ but not libc++.
|
||||
2. GGL currently requires templated kernel fucntions passed to hipLaunchKernel to be specialized. e.g.:
|
||||
|
||||
```
|
||||
hipLaunchKernel(vector_square<float>, dim3(blocks), dim3(threadsPerBlock), 0, nullptr, C_d, A_d, N);
|
||||
```
|
||||
GGL was only tested with Ubuntu16.04, assuming HCC and HIP only link against libstdc++ but not libc++.
|
||||
|
||||
@@ -323,7 +323,6 @@ namespace glo_tests
|
||||
make_lambda_wrapper(kernel_name, __VA_ARGS__), \
|
||||
##__VA_ARGS__); \
|
||||
}
|
||||
#warning "GGL hipLaunchKernel Reached"
|
||||
#define hipLaunchKernel( \
|
||||
kernel_name, \
|
||||
num_blocks, \
|
||||
|
||||
@@ -32,7 +32,7 @@ THE SOFTWARE.
|
||||
#include <iostream>
|
||||
|
||||
#ifndef GENERIC_GRID_LAUNCH
|
||||
#define GENERIC_GRID_LAUNCH 0
|
||||
#define GENERIC_GRID_LAUNCH 0
|
||||
#endif
|
||||
|
||||
#include <hip/hcc_detail/host_defines.h>
|
||||
|
||||
@@ -32,7 +32,7 @@ THE SOFTWARE.
|
||||
|
||||
// Add guard to Generic Grid Launch method
|
||||
#ifndef GENERIC_GRID_LAUNCH
|
||||
#define GENERIC_GRID_LAUNCH 0
|
||||
#define GENERIC_GRID_LAUNCH 0
|
||||
#endif
|
||||
|
||||
#ifdef __HCC__
|
||||
|
||||
@@ -83,7 +83,7 @@ int main(int argc, char *argv[])
|
||||
const unsigned threadsPerBlock = 256;
|
||||
|
||||
printf ("info: launch 'vector_square' kernel\n");
|
||||
hipLaunchKernel(vector_square<float>, dim3(blocks), dim3(threadsPerBlock), 0, nullptr, C_d, A_d, N);
|
||||
hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, nullptr, C_d, A_d, N);
|
||||
|
||||
printf ("info: copy Device2Host\n");
|
||||
CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
Referência em uma Nova Issue
Bloquear um usuário