AI/rocm-systems
2
0
Rozštěpit 0
Zdrojový kód Úkoly Pull requesty Akce Balíčky Projekty Vydání Wiki Aktivita

Merge 'master' into 'amd-master'

Procházet zdrojové kódy 
Change-Id: I99fcb0d69a58ce09065991a580662660d2941f9f


[ROCm/clr commit: 0e1ef39580]
Tento commit je obsažen v:
Jenkins
2019-09-12 06:06:14 -04:00
rodič 16adcbf9bf ae66e0b3a1
revize b834dd2c35
6 změnil soubory, kde provedl 102 přidání a 71 odebrání
Stáhněte soubor záplaty Stáhněte rozdílový soubor Rozbalit všechny soubory Sbalit všechny soubory
projects/clr/hipamd/bin/hipify-perl
+31 -12
Zobrazit soubor
@@ -1662,18 +1662,37 @@ while (@ARGV) {
# match uses ? for <.*> which will be unitialized if this is not present in launch syntax. # match uses ? for <.*> which will be unitialized if this is not present in launch syntax.
no warnings qw/uninitialized/; no warnings qw/uninitialized/;
my $k = 0; my $k = 0;
# Handle the <<numBlocks, blockDim, sharedSize, stream>>> syntax with empty args:
$k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), $5, $6)/g; # Handle the kern<...><<<numBlocks, blockDim, sharedSize, stream>>> syntax with empty args:
# Handle the <<numBlocks, blockDim, sharedSize, stream>>> syntax with non-empty args: $k += s/(\w+)\s*<(.+)>\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), $5, $6)/g;
$k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), $5, $6, /g; # Handle the kern<<<numBlocks, blockDim, sharedSize, stream>>> syntax with empty args:
# Handle the <<numBlocks, blockDim, sharedSize>>> syntax with empty args: $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, $5)/g;
$k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), $5, 0)/g;
# Handle the <<numBlocks, blockDim, sharedSize>>> syntax with non-empty args: # Handle the kern<...><<<numBlocks, blockDim, sharedSize, stream>>> syntax with non-empty args:
$k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), $5, 0, /g; $k += s/(\w+)\s*<(.+)>\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), $5, $6, /g;
# Handle the <<numBlocks, blockDim>>> syntax with empty args: # Handle the kern<<<numBlocks, blockDim, sharedSize, stream>>> syntax with non-empty args:
$k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), 0, 0)/g; $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, $5, /g;
# Handle the <<numBlocks, blockDim>>> syntax with non-empty args:
$k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), 0, 0, /g; # Handle the kern<...><<<numBlocks, blockDim, sharedSize>>> syntax with empty args:
$k += s/(\w+)\s*<(.+)>\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), $5, 0)/g;
# Handle the kern<<<numBlocks, blockDim, sharedSize>>> syntax with empty args:
$k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, 0)/g;
# Handle the kern<...><<<numBlocks, blockDim, sharedSize>>> syntax with non-empty args:
$k += s/(\w+)\s*<(.+)>\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), $5, 0, /g;
# Handle the kern<<<numBlocks, blockDim, sharedSize>>> syntax with non-empty args:
$k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, 0, /g;
# Handle the kern<...><<<numBlocks, blockDim>>> syntax with empty args:
$k += s/(\w+)\s*<(.+)>\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), 0, 0)/g;
# Handle the kern<<<numBlocks, blockDim>>> syntax with empty args:
$k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0)/g;
# Handle the kern<...><<<numBlocks, blockDim>>> syntax with non-empty args:
$k += s/(\w+)\s*<(.+)>\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), 0, 0, /g;
# Handle the kern<<<numBlocks, blockDim>>> syntax with non-empty args:
$k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0, /g;
if ($k) { if ($k) {
$ft{'kern_launch'} += $k; $ft{'kern_launch'} += $k;
$Tkernels{$1} ++; $Tkernels{$1} ++;
projects/clr/hipamd/hipify-clang/src/CUDA2HIP.cpp
+37 -36
Zobrazit soubor
@@ -25,47 +25,48 @@ THE SOFTWARE.
// Maps CUDA header names to HIP header names // Maps CUDA header names to HIP header names
const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP{ const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP{
// CUDA includes // CUDA includes
{"cuda.h", {"hip/hip_runtime.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_DRIVER}}, {"cuda.h", {"hip/hip_runtime.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_DRIVER}},
{"cuda_runtime.h", {"hip/hip_runtime.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_RUNTIME}}, {"cuda_runtime.h", {"hip/hip_runtime.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_RUNTIME}},
{"cuda_runtime_api.h", {"hip/hip_runtime_api.h", "", CONV_INCLUDE, API_RUNTIME}}, {"cuda_runtime_api.h", {"hip/hip_runtime_api.h", "", CONV_INCLUDE, API_RUNTIME}},
{"channel_descriptor.h", {"hip/channel_descriptor.h", "", CONV_INCLUDE, API_RUNTIME}}, {"channel_descriptor.h", {"hip/channel_descriptor.h", "", CONV_INCLUDE, API_RUNTIME}},
{"device_functions.h", {"hip/device_functions.h", "", CONV_INCLUDE, API_RUNTIME}}, {"device_functions.h", {"hip/device_functions.h", "", CONV_INCLUDE, API_RUNTIME}},
{"driver_types.h", {"hip/driver_types.h", "", CONV_INCLUDE, API_RUNTIME}}, {"driver_types.h", {"hip/driver_types.h", "", CONV_INCLUDE, API_RUNTIME}},
{"cuda_fp16.h", {"hip/hip_fp16.h", "", CONV_INCLUDE, API_RUNTIME}}, {"cuda_fp16.h", {"hip/hip_fp16.h", "", CONV_INCLUDE, API_RUNTIME}},
{"cuda_texture_types.h", {"hip/hip_texture_types.h", "", CONV_INCLUDE, API_RUNTIME}}, {"cuda_texture_types.h", {"hip/hip_texture_types.h", "", CONV_INCLUDE, API_RUNTIME}},
{"texture_fetch_functions.h", {"", "", CONV_INCLUDE, API_RUNTIME}}, {"texture_fetch_functions.h", {"", "", CONV_INCLUDE, API_RUNTIME}},
{"vector_types.h", {"hip/hip_vector_types.h", "", CONV_INCLUDE, API_RUNTIME}}, {"vector_types.h", {"hip/hip_vector_types.h", "", CONV_INCLUDE, API_RUNTIME}},
{"cuda_profiler_api.h", {"hip/hip_profile.h", "", CONV_INCLUDE, API_RUNTIME}}, {"cuda_profiler_api.h", {"hip/hip_profile.h", "", CONV_INCLUDE, API_RUNTIME}},
{"cooperative_groups.h", {"hip/hip_cooperative_groups.h", "", CONV_INCLUDE, API_RUNTIME}},
// cuComplex includes // cuComplex includes
{"cuComplex.h", {"hip/hip_complex.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_COMPLEX}}, {"cuComplex.h", {"hip/hip_complex.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_COMPLEX}},
// cuBLAS includes // cuBLAS includes
{"cublas.h", {"hipblas.h", "rocblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}}, {"cublas.h", {"hipblas.h", "rocblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}},
{"cublas_v2.h", {"hipblas.h", "rocblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}}, {"cublas_v2.h", {"hipblas.h", "rocblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}},
// cuRAND includes // cuRAND includes
{"curand.h", {"hiprand.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_RAND}}, {"curand.h", {"hiprand.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_RAND}},
{"curand_kernel.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_kernel.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_discrete.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_discrete.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_discrete2.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_discrete2.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_globals.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_globals.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_lognormal.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_lognormal.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_mrg32k3a.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_mrg32k3a.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_mtgp32.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_mtgp32.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_mtgp32_host.h", {"hiprand_mtgp32_host.h", "", CONV_INCLUDE, API_RAND}}, {"curand_mtgp32_host.h", {"hiprand_mtgp32_host.h", "", CONV_INCLUDE, API_RAND}},
{"curand_mtgp32_kernel.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_mtgp32_kernel.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_mtgp32dc_p_11213.h", {"rocrand_mtgp32_11213.h", "", CONV_INCLUDE, API_RAND}}, {"curand_mtgp32dc_p_11213.h", {"rocrand_mtgp32_11213.h", "", CONV_INCLUDE, API_RAND}},
{"curand_normal.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_normal.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_normal_static.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_normal_static.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_philox4x32_x.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_philox4x32_x.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_poisson.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_poisson.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_precalc.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_precalc.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_uniform.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_uniform.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
// cuDNN includes // cuDNN includes
{"cudnn.h", {"hipDNN.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_DNN}}, {"cudnn.h", {"hipDNN.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_DNN}},
// cuFFT includes // cuFFT includes
{"cufft.h", {"hipfft.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_FFT}}, {"cufft.h", {"hipfft.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_FFT}},
// cuBLAS includes // cuSPARSE includes
{"cusparse.h", {"hipsparse.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_SPARSE}}, {"cusparse.h", {"hipsparse.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_SPARSE}},
{"cusparse_v2.h", {"hipsparse.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_SPARSE}}, {"cusparse_v2.h", {"hipsparse.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_SPARSE}},
// CAFFE2 includes // CAFFE2 includes
{"caffe2/core/common_gpu.h", {"caffe2/core/hip/common_gpu.h", "", CONV_INCLUDE, API_CAFFE2, UNSUPPORTED}}, {"caffe2/core/common_gpu.h", {"caffe2/core/hip/common_gpu.h", "", CONV_INCLUDE, API_CAFFE2, UNSUPPORTED}},
{"caffe2/core/context_gpu.h", {"caffe2/core/hip/context_gpu.h", "", CONV_INCLUDE, API_CAFFE2, UNSUPPORTED}}, {"caffe2/core/context_gpu.h", {"caffe2/core/hip/context_gpu.h", "", CONV_INCLUDE, API_CAFFE2, UNSUPPORTED}},
projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp
+21 -10
Zobrazit soubor
@@ -289,24 +289,35 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc
if (!launchKernel) { if (!launchKernel) {
return false; return false;
} }
const clang::Expr* calleeExpr = launchKernel->getCallee();
if (!calleeExpr) {
return false;
}
const clang::FunctionDecl *caleeDecl = launchKernel->getDirectCallee();
if (!caleeDecl) {
return false;
}
const clang::CallExpr* config = launchKernel->getConfig();
if (!config) {
return false;
}
clang::SmallString<40> XStr; clang::SmallString<40> XStr;
llvm::raw_svector_ostream OS(XStr); llvm::raw_svector_ostream OS(XStr);
clang::LangOptions DefaultLangOptions; clang::LangOptions DefaultLangOptions;
clang::SourceManager* SM = Result.SourceManager; clang::SourceManager* SM = Result.SourceManager;
OS << "hipLaunchKernelGGL(";
const clang::Expr& calleeExpr = *(launchKernel->getCallee()); if (caleeDecl->isTemplateInstantiation()) OS << "(";
OS << "hipLaunchKernelGGL(" << readSourceText(*SM, calleeExpr.getSourceRange()) << ", "; OS << readSourceText(*SM, calleeExpr->getSourceRange());
if (caleeDecl->isTemplateInstantiation()) OS << ")";
OS << ", ";
// Next up are the four kernel configuration parameters, the last two of which are optional and default to zero. // Next up are the four kernel configuration parameters, the last two of which are optional and default to zero.
const clang::CallExpr& config = *(launchKernel->getConfig());
// Copy the two dimensional arguments verbatim. // Copy the two dimensional arguments verbatim.
OS << "dim3(" << readSourceText(*SM, config.getArg(0)->getSourceRange()) << "), "; OS << "dim3(" << readSourceText(*SM, config->getArg(0)->getSourceRange()) << "), ";
OS << "dim3(" << readSourceText(*SM, config.getArg(1)->getSourceRange()) << "), "; OS << "dim3(" << readSourceText(*SM, config->getArg(1)->getSourceRange()) << "), ";
// The stream/memory arguments default to zero if omitted. // The stream/memory arguments default to zero if omitted.
OS << stringifyZeroDefaultedArg(*SM, config.getArg(2)) << ", "; OS << stringifyZeroDefaultedArg(*SM, config->getArg(2)) << ", ";
OS << stringifyZeroDefaultedArg(*SM, config.getArg(3)); OS << stringifyZeroDefaultedArg(*SM, config->getArg(3));
// If there are ordinary arguments to the kernel, just copy them verbatim into our new call. // If there are ordinary arguments to the kernel, just copy them verbatim into our new call.
int numArgs = launchKernel->getNumArgs(); int numArgs = launchKernel->getNumArgs();
projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp
+5 -5
Zobrazit soubor
@@ -121,7 +121,7 @@ struct runner
const size_t states_size = blocks * threads; const size_t states_size = blocks * threads;
// CHECK: CUDA_CALL(hipMalloc((void **)&states, states_size * sizeof(GeneratorState))); // CHECK: CUDA_CALL(hipMalloc((void **)&states, states_size * sizeof(GeneratorState)));
CUDA_CALL(cudaMalloc((void **)&states, states_size * sizeof(GeneratorState))); CUDA_CALL(cudaMalloc((void **)&states, states_size * sizeof(GeneratorState)));
// CHECK: hipLaunchKernelGGL(init_kernel, dim3(blocks), dim3(threads), 0, 0, states, seed, offset); // CHECK: hipLaunchKernelGGL((init_kernel), dim3(blocks), dim3(threads), 0, 0, states, seed, offset);
init_kernel<<<blocks, threads>>>(states, seed, offset); init_kernel<<<blocks, threads>>>(states, seed, offset);
// CHECK: CUDA_CALL(hipPeekAtLastError()); // CHECK: CUDA_CALL(hipPeekAtLastError());
// CHECK: CUDA_CALL(hipDeviceSynchronize()); // CHECK: CUDA_CALL(hipDeviceSynchronize());
@@ -142,7 +142,7 @@ struct runner
const GenerateFunc& generate_func, const GenerateFunc& generate_func,
const Extra extra) const Extra extra)
{ {
// CHECK: hipLaunchKernelGGL(generate_kernel, dim3(blocks), dim3(threads), 0, 0, states, data, size, generate_func, extra); // CHECK: hipLaunchKernelGGL((generate_kernel), dim3(blocks), dim3(threads), 0, 0, states, data, size, generate_func, extra);
generate_kernel<<<blocks, threads>>>(states, data, size, generate_func, extra); generate_kernel<<<blocks, threads>>>(states, data, size, generate_func, extra);
} }
}; };
@@ -223,7 +223,7 @@ struct runner<curandStateMtgp32_t>
const GenerateFunc& generate_func, const GenerateFunc& generate_func,
const Extra extra) const Extra extra)
{ {
// CHECK: hipLaunchKernelGGL(generate_kernel, dim3(std::min((size_t)200, blocks)), dim3(256), 0, 0, states, data, size, generate_func, extra); // CHECK: hipLaunchKernelGGL((generate_kernel), dim3(std::min((size_t)200, blocks)), dim3(256), 0, 0, states, data, size, generate_func, extra);
generate_kernel<<<std::min((size_t)200, blocks), 256>>>(states, data, size, generate_func, extra); generate_kernel<<<std::min((size_t)200, blocks), 256>>>(states, data, size, generate_func, extra);
} }
}; };
@@ -304,7 +304,7 @@ struct runner<curandStateSobol32_t>
CUDA_CALL(cudaMemcpy(directions, h_directions, size, cudaMemcpyHostToDevice)); CUDA_CALL(cudaMemcpy(directions, h_directions, size, cudaMemcpyHostToDevice));
const size_t blocks_x = next_power2((blocks + dimensions - 1) / dimensions); const size_t blocks_x = next_power2((blocks + dimensions - 1) / dimensions);
// CHECK: hipLaunchKernelGGL(init_kernel, dim3(dim3(blocks_x, dimensions)), dim3(threads), 0, 0, states, directions, offset); // CHECK: hipLaunchKernelGGL((init_kernel), dim3(dim3(blocks_x, dimensions)), dim3(threads), 0, 0, states, directions, offset);
init_kernel<<<dim3(blocks_x, dimensions), threads>>>(states, directions, offset); init_kernel<<<dim3(blocks_x, dimensions), threads>>>(states, directions, offset);
// CHECK: CUDA_CALL(hipPeekAtLastError()); // CHECK: CUDA_CALL(hipPeekAtLastError());
// CHECK: CUDA_CALL(hipDeviceSynchronize()); // CHECK: CUDA_CALL(hipDeviceSynchronize());
@@ -329,7 +329,7 @@ struct runner<curandStateSobol32_t>
const Extra extra) const Extra extra)
{ {
const size_t blocks_x = next_power2((blocks + dimensions - 1) / dimensions); const size_t blocks_x = next_power2((blocks + dimensions - 1) / dimensions);
// CHECK: hipLaunchKernelGGL(generate_kernel, dim3(dim3(blocks_x, dimensions)), dim3(threads), 0, 0, states, data, size / dimensions, generate_func, extra); // CHECK: hipLaunchKernelGGL((generate_kernel), dim3(dim3(blocks_x, dimensions)), dim3(threads), 0, 0, states, data, size / dimensions, generate_func, extra);
generate_kernel<<<dim3(blocks_x, dimensions), threads>>>(states, data, size / dimensions, generate_func, extra); generate_kernel<<<dim3(blocks_x, dimensions), threads>>>(states, data, size / dimensions, generate_func, extra);
} }
}; };
projects/clr/hipamd/tests/hipify-clang/unit_tests/samples/axpy.cu
+7 -7
Zobrazit soubor
@@ -10,7 +10,7 @@
#define KERNEL_CALL_AS_MACRO axpy<float><<<1, kDataLen>>> #define KERNEL_CALL_AS_MACRO axpy<float><<<1, kDataLen>>>
#define KERNEL_NAME_MACRO axpy<float> #define KERNEL_NAME_MACRO axpy<float>
// CHECK: #define COMPLETE_LAUNCH hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y) // CHECK: #define COMPLETE_LAUNCH hipLaunchKernelGGL((axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y)
#define COMPLETE_LAUNCH axpy<<<1, kDataLen>>>(a, device_x, device_y) #define COMPLETE_LAUNCH axpy<<<1, kDataLen>>>(a, device_x, device_y)
@@ -48,22 +48,22 @@ int main(int argc, char* argv[]) {
cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), cudaMemcpyHostToDevice);
// Launch the kernel in numerous different strange ways to exercise the prerocessor. // Launch the kernel in numerous different strange ways to exercise the prerocessor.
// CHECK: hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); // CHECK: hipLaunchKernelGGL((axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
axpy<<<1, kDataLen>>>(a, device_x, device_y); axpy<<<1, kDataLen>>>(a, device_x, device_y);
// CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); // CHECK: hipLaunchKernelGGL((axpy<float>), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
axpy<float><<<1, kDataLen>>>(a, device_x, device_y); axpy<float><<<1, kDataLen>>>(a, device_x, device_y);
// CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, a, TOKEN_PASTE(device, _x), device_y); // CHECK: hipLaunchKernelGGL((axpy<float>), dim3(1), dim3(kDataLen), 0, 0, a, TOKEN_PASTE(device, _x), device_y);
axpy<float><<<1, kDataLen>>>(a, TOKEN_PASTE(device, _x), device_y); axpy<float><<<1, kDataLen>>>(a, TOKEN_PASTE(device, _x), device_y);
// CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO); // CHECK: hipLaunchKernelGGL((axpy<float>), dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
axpy<float><<<1, kDataLen>>>(ARG_LIST_AS_MACRO); axpy<float><<<1, kDataLen>>>(ARG_LIST_AS_MACRO);
// CHECK: hipLaunchKernelGGL(KERNEL_NAME_MACRO, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO); // CHECK: hipLaunchKernelGGL((KERNEL_NAME_MACRO), dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
KERNEL_NAME_MACRO<<<1, kDataLen>>>(ARG_LIST_AS_MACRO); KERNEL_NAME_MACRO<<<1, kDataLen>>>(ARG_LIST_AS_MACRO);
// CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO); // CHECK: hipLaunchKernelGGL((axpy<float>), dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
KERNEL_CALL_AS_MACRO(ARG_LIST_AS_MACRO); KERNEL_CALL_AS_MACRO(ARG_LIST_AS_MACRO);
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), 0, 0); // CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), 0, 0);
projects/clr/hipamd/tests/hipify-clang/unit_tests/samples/square.cu
+1 -1
Zobrazit soubor
@@ -94,7 +94,7 @@ int main(int argc, char *argv[])
const unsigned threadsPerBlock = 256; const unsigned threadsPerBlock = 256;
printf ("info: launch 'vector_square' kernel\n"); printf ("info: launch 'vector_square' kernel\n");
// CHECK: hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N); // CHECK: hipLaunchKernelGGL((vector_square), dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
vector_square <<<blocks, threadsPerBlock>>> (C_d, A_d, N); vector_square <<<blocks, threadsPerBlock>>> (C_d, A_d, N);
printf ("info: copy Device2Host\n"); printf ("info: copy Device2Host\n");