diff --git a/languages/clang/duchain/parsesession.cpp b/languages/clang/duchain/parsesession.cpp --- a/languages/clang/duchain/parsesession.cpp +++ b/languages/clang/duchain/parsesession.cpp @@ -78,6 +78,12 @@ return {QByteArrayLiteral("-xcl")}; } + // TODO: No proper mime type detection possible yet + // cf. https://bugs.freedesktop.org/show_bug.cgi?id=23700 + if (path.endsWith(QLatin1String(".cu"), Qt::CaseInsensitive)) { + return {QByteArrayLiteral("-xcuda")}; + } + if (parserSettings.parserOptions.isEmpty()) { // The parserOptions can be empty for some unit tests that use ParseSession directly auto defaultArguments = ClangSettingsManager::self()->parserSettings(path).toClangAPI(); diff --git a/languages/clang/kdevclangsupport.json b/languages/clang/kdevclangsupport.json --- a/languages/clang/kdevclangsupport.json +++ b/languages/clang/kdevclangsupport.json @@ -1,72 +1,74 @@ { "KPlugin": { - "Description": "C/C++ Language Support (Clang-based)", - "Description[ca@valencia]": "Implementació del llenguatge C/C++ (basat en Clang)", - "Description[ca]": "Implementació del llenguatge C/C++ (basat en Clang)", - "Description[cs]": "Podpora jazyka C/C++ (založeno na Clang)", - "Description[de]": "Sprachunterstützung für C/C++ (Clang-basiert)", - "Description[es]": "Implementación del lenguaje C/C++ (basado en Clang)", - "Description[fi]": "C/C++-kielituki (Clang-pohjainen)", - "Description[fr]": "Prise en charge du langage C/C++ (utilisant Clang)", - "Description[gl]": "Compatibilidade con C e C++ (baseada en Clang)", - "Description[it]": "Supporto al linguaggio C/C++ (basato su Clang)", - "Description[nl]": "Ondersteuning voor de talen C/C++ (Clang-based)", - "Description[pl]": "Obsługa języków C/C++ (na podstawie Clang)", - "Description[pt]": "Suporte para a Linguagem C/C++ (baseado no Clang)", - "Description[pt_BR]": "Suporte à linguagem C/C++ (baseado no Clang)", - "Description[sk]": "Podpora jazyka C/C++ (založená na triedach)", - "Description[sl]": "Podpora jeziku C/C++ (temelječa na Clang)", - "Description[sv]": "C/C++ språkstöd (Clang-baserat)", - "Description[uk]": "Підтримка мови C/C++ на основі Clang", - "Description[x-test]": "xxC/C++ Language Support (Clang-based)xx", - "Description[zh_CN]": "C/C++ 语言支持 (基于 Clang)", + "Description": "C/C++ Language Support (Clang-based)", + "Description[ca@valencia]": "Implementació del llenguatge C/C++ (basat en Clang)", + "Description[ca]": "Implementació del llenguatge C/C++ (basat en Clang)", + "Description[cs]": "Podpora jazyka C/C++ (založeno na Clang)", + "Description[de]": "Sprachunterstützung für C/C++ (Clang-basiert)", + "Description[es]": "Implementación del lenguaje C/C++ (basado en Clang)", + "Description[fi]": "C/C++-kielituki (Clang-pohjainen)", + "Description[fr]": "Prise en charge du langage C/C++ (utilisant Clang)", + "Description[gl]": "Compatibilidade con C e C++ (baseada en Clang)", + "Description[it]": "Supporto al linguaggio C/C++ (basato su Clang)", + "Description[nl]": "Ondersteuning voor de talen C/C++ (Clang-based)", + "Description[pl]": "Obsługa języków C/C++ (na podstawie Clang)", + "Description[pt]": "Suporte para a Linguagem C/C++ (baseado no Clang)", + "Description[pt_BR]": "Suporte à linguagem C/C++ (baseado no Clang)", + "Description[sk]": "Podpora jazyka C/C++ (založená na triedach)", + "Description[sl]": "Podpora jeziku C/C++ (temelječa na Clang)", + "Description[sv]": "C/C++ språkstöd (Clang-baserat)", + "Description[uk]": "Підтримка мови C/C++ на основі Clang", + "Description[x-test]": "xxC/C++ Language Support (Clang-based)xx", + "Description[zh_CN]": "C/C++ 语言支持 (基于 Clang)", "Icon": "text-x-c++src", - "Id": "kdevclangsupport", - "Name": "C++ Support", - "Name[bs]": "C++ Podrška", - "Name[ca@valencia]": "Implementació del C++", - "Name[ca]": "Implementació del C++", - "Name[cs]": "Podpora C++", - "Name[de]": "Unterstützung für C++", - "Name[es]": "Implementación de C++", - "Name[fi]": "C++-tuki", - "Name[fr]": "Prise en charge du C++", - "Name[gl]": "Compatibilidade con C++", - "Name[hu]": "C++ támogatás", - "Name[it]": "Supporto C++", - "Name[nl]": "Ondersteuning voor C++", - "Name[pl]": "Obsługa C++", - "Name[pt]": "Suporte a C/C++", - "Name[pt_BR]": "Suporte à C++", - "Name[ru]": "Поддержка C++", - "Name[sk]": "Podpora C++", - "Name[sl]": "Podpora za C++", - "Name[sv]": "C++ stöd", - "Name[tr]": "C++ Desteği", - "Name[uk]": "Підтримка C++", - "Name[x-test]": "xxC++ Supportxx", - "Name[zh_CN]": "C++ 支持", + "Id": "kdevclangsupport", + "Name": "C++ Support", + "Name[bs]": "C++ Podrška", + "Name[ca@valencia]": "Implementació del C++", + "Name[ca]": "Implementació del C++", + "Name[cs]": "Podpora C++", + "Name[de]": "Unterstützung für C++", + "Name[es]": "Implementación de C++", + "Name[fi]": "C++-tuki", + "Name[fr]": "Prise en charge du C++", + "Name[gl]": "Compatibilidade con C++", + "Name[hu]": "C++ támogatás", + "Name[it]": "Supporto C++", + "Name[nl]": "Ondersteuning voor C++", + "Name[pl]": "Obsługa C++", + "Name[pt]": "Suporte a C/C++", + "Name[pt_BR]": "Suporte à C++", + "Name[ru]": "Поддержка C++", + "Name[sk]": "Podpora C++", + "Name[sl]": "Podpora za C++", + "Name[sv]": "C++ stöd", + "Name[tr]": "C++ Desteği", + "Name[uk]": "Підтримка C++", + "Name[x-test]": "xxC++ Supportxx", + "Name[zh_CN]": "C++ 支持", "ServiceTypes": [ "KDevelop/Plugin" ] - }, + }, "X-KDevelop-Interfaces": [ "ILanguageSupport" - ], + ], "X-KDevelop-Languages": [ - "C", - "C++", - "OpenCL", + "C", + "C++", + "OpenCL", + "CUDA C", "Objective-C" - ], - "X-KDevelop-LoadMode": "AlwaysOn", - "X-KDevelop-Mode": "NoGUI", + ], + "X-KDevelop-LoadMode": "AlwaysOn", + "X-KDevelop-Mode": "NoGUI", "X-KDevelop-SupportedMimeTypes": [ - "text/x-chdr", - "text/x-c++hdr", - "text/x-csrc", - "text/x-c++src", - "text/x-opencl-src", + "text/x-chdr", + "text/x-c++hdr", + "text/x-csrc", + "text/x-c++src", + "text/x-openclsrc", + "text/vnd.nvidia.cuda.csrc", "text/x-objcsrc" ] } diff --git a/languages/clang/tests/files/cuda.cu b/languages/clang/tests/files/cuda.cu new file mode 100644 --- /dev/null +++ b/languages/clang/tests/files/cuda.cu @@ -0,0 +1,144 @@ +// Thread block size +#define BLOCK_SIZE 16 + +// Matrices are stored in row-major order: +// M(row, col) = *(M.elements + row * M.stride + col) +typedef struct { + int width; + int height; + int stride; + float* elements; +} Matrix; + +// Get a matrix element +__device__ float GetElement(const Matrix A, int row, int col) +{ + return A.elements[row * A.stride + col]; +} + +// Set a matrix element +__device__ void SetElement(Matrix A, int row, int col, + float value) +{ + A.elements[row * A.stride + col] = value; +} + +// Get the BLOCK_SIZExBLOCK_SIZE sub-matrix Asub of A that is +// located col sub-matrices to the right and row sub-matrices down +// from the upper-left corner of A +__device__ Matrix GetSubMatrix(Matrix A, int row, int col) +{ + Matrix Asub; + Asub.width = BLOCK_SIZE; + Asub.height = BLOCK_SIZE; + Asub.stride = A.stride; + Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row + + BLOCK_SIZE * col]; + return Asub; +} + + +__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) +{ + // Block row and column + int blockRow = blockIdx.y; + int blockCol = blockIdx.x; + + // Each thread block computes one sub-matrix Csub of C + Matrix Csub = GetSubMatrix(C, blockRow, blockCol); + + // Each thread computes one element of Csub + // by accumulating results into Cvalue + float Cvalue = 0; + + // Thread row and column within Csub + int row = threadIdx.y; + int col = threadIdx.x; + + // Loop over all the sub-matrices of A and B that are + // required to compute Csub + // Multiply each pair of sub-matrices together + // and accumulate the results + for (int m = 0; m < (A.width / BLOCK_SIZE); ++m) { + + // Get sub-matrix Asub of A + Matrix Asub = GetSubMatrix(A, blockRow, m); + + // Get sub-matrix Bsub of B + Matrix Bsub = GetSubMatrix(B, m, blockCol); + + // Shared memory used to store Asub and Bsub respectively + __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; + __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; + + // Load Asub and Bsub from device memory to shared memory + // Each thread loads one element of each sub-matrix + As[row][col] = GetElement(Asub, row, col); + Bs[row][col] = GetElement(Bsub, row, col); + + // Synchronize to make sure the sub-matrices are loaded + // before starting the computation + __syncthreads(); + + // Multiply Asub and Bsub together + for (int e = 0; e < BLOCK_SIZE; ++e) + Cvalue += As[row][e] * Bs[e][col]; + + // Synchronize to make sure that the preceding + // computation is done before loading two new + // sub-matrices of A and B in the next iteration + __syncthreads(); + } + + // Write Csub to device memory + // Each thread writes one element + SetElement(Csub, row, col, Cvalue); +} + +// Matrix multiplication - Host code +// Matrix dimensions are assumed to be multiples of BLOCK_SIZE +void MatMul(const Matrix A, const Matrix B, Matrix C) +{ + // Load A and B to device memory + Matrix d_A; + d_A.width = d_A.stride = A.width; d_A.height = A.height; + size_t size = A.width * A.height * sizeof(float); + cudaMalloc(&d_A.elements, size); + cudaMemcpy(d_A.elements, A.elements, size, + cudaMemcpyHostToDevice); + Matrix d_B; + d_B.width = d_B.stride = B.width; d_B.height = B.height; + size = B.width * B.height * sizeof(float); + + cudaMalloc(&d_B.elements, size); + cudaMemcpy(d_B.elements, B.elements, size, + cudaMemcpyHostToDevice); + + // Allocate C in device memory + Matrix d_C; + d_C.width = d_C.stride = C.width; d_C.height = C.height; + size = C.width * C.height * sizeof(float); + cudaMalloc(&d_C.elements, size); + + // Invoke kernel + dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); + dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y); + MatMulKernel<<>>(d_A, d_B, d_C); + + // Read C from device memory + cudaMemcpy(C.elements, d_C.elements, size, + cudaMemcpyDeviceToHost); + + // Free device memory + cudaFree(d_A.elements); + cudaFree(d_B.elements); + cudaFree(d_C.elements); +} + +// Matrix multiplication kernel called by MatMul() + +int main() +{ + Matrix a, b, c; + MatMul(a, b, c); +} diff --git a/languages/clang/tests/test_files.cpp b/languages/clang/tests/test_files.cpp --- a/languages/clang/tests/test_files.cpp +++ b/languages/clang/tests/test_files.cpp @@ -67,7 +67,7 @@ { QTest::addColumn("fileName"); const QString testDirPath = TEST_FILES_DIR; - const QStringList files = QDir(testDirPath).entryList({"*.h", "*.cpp", "*.c", "*.cl"}, QDir::Files); + const QStringList files = QDir(testDirPath).entryList({"*.h", "*.cpp", "*.c", "*.cl", "*.cu"}, QDir::Files); foreach (const QString& file, files) { QTest::newRow(file.toUtf8().constData()) << QString(testDirPath + '/' + file); } diff --git a/languages/plugins/custom-definesandincludes/compilerprovider/settingsmanager.h b/languages/plugins/custom-definesandincludes/compilerprovider/settingsmanager.h --- a/languages/plugins/custom-definesandincludes/compilerprovider/settingsmanager.h +++ b/languages/plugins/custom-definesandincludes/compilerprovider/settingsmanager.h @@ -41,6 +41,7 @@ QString cArguments; QString cppArguments; QString openClArguments; + QString cudaArguments; bool parseAmbiguousAsCPP; }; @@ -67,6 +68,7 @@ C, Cpp, OpenCl, + Cuda, ObjC, Other = 100 diff --git a/languages/plugins/custom-definesandincludes/compilerprovider/settingsmanager.cpp b/languages/plugins/custom-definesandincludes/compilerprovider/settingsmanager.cpp --- a/languages/plugins/custom-definesandincludes/compilerprovider/settingsmanager.cpp +++ b/languages/plugins/custom-definesandincludes/compilerprovider/settingsmanager.cpp @@ -85,7 +85,8 @@ const static ParserArguments arguments{ QStringLiteral("-ferror-limit=100 -fspell-checking -Wdocumentation -Wunused-parameter -Wunreachable-code -Wall -std=c99"), QStringLiteral("-ferror-limit=100 -fspell-checking -Wdocumentation -Wunused-parameter -Wunreachable-code -Wall -std=c++11"), - QStringLiteral("-ferror-limit=100 -fspell-checking -Wdocumentation -Wunused-parameter -Wunreachable-code -Wall"), + QStringLiteral("-ferror-limit=100 -fspell-checking -Wdocumentation -Wunused-parameter -Wunreachable-code -Wall -std=c++11"), + QStringLiteral("-ferror-limit=100 -fspell-checking -Wdocumentation -Wunused-parameter -Wunreachable-code -Wall -std=c++11"), true }; @@ -405,6 +406,12 @@ return OpenCl; } + // TODO: No proper mime type detection possible yet + // cf. https://bugs.freedesktop.org/show_bug.cgi?id=23700 + if (path.lastPathSegment().endsWith(QLatin1String(".cu"), Qt::CaseInsensitive)) { + return Cuda; + } + return C; } diff --git a/languages/plugins/custom-definesandincludes/definesandincludesmanager.cpp b/languages/plugins/custom-definesandincludes/definesandincludesmanager.cpp --- a/languages/plugins/custom-definesandincludes/definesandincludesmanager.cpp +++ b/languages/plugins/custom-definesandincludes/definesandincludesmanager.cpp @@ -109,6 +109,8 @@ return arguments.cppArguments; case Utils::OpenCl: return arguments.openClArguments; + case Utils::Cuda: + return arguments.cudaArguments; case Utils::ObjC: return QString(); case Utils::Other: diff --git a/languages/plugins/custom-definesandincludes/kcm_widget/parserwidget.h b/languages/plugins/custom-definesandincludes/kcm_widget/parserwidget.h --- a/languages/plugins/custom-definesandincludes/kcm_widget/parserwidget.h +++ b/languages/plugins/custom-definesandincludes/kcm_widget/parserwidget.h @@ -56,6 +56,8 @@ void textEdited(); void languageStandardChangedC(const QString& standard); void languageStandardChangedCpp(const QString& standard); + void languageStandardChangedOpenCl(const QString& standard); + void languageStandardChangedCuda(const QString& standard); void updateEnablements(); private: diff --git a/languages/plugins/custom-definesandincludes/kcm_widget/parserwidget.cpp b/languages/plugins/custom-definesandincludes/kcm_widget/parserwidget.cpp --- a/languages/plugins/custom-definesandincludes/kcm_widget/parserwidget.cpp +++ b/languages/plugins/custom-definesandincludes/kcm_widget/parserwidget.cpp @@ -71,13 +71,17 @@ connect(m_ui->parserOptionsC, &QLineEdit::textEdited, this, &ParserWidget::textEdited); connect(m_ui->parserOptionsCpp, &QLineEdit::textEdited, this, &ParserWidget::textEdited); + connect(m_ui->parserOptionsOpenCl, &QLineEdit::textEdited, this, &ParserWidget::textEdited); + connect(m_ui->parserOptionsCuda, &QLineEdit::textEdited, this, &ParserWidget::textEdited); connect(m_ui->parseHeadersInPlainC, &QCheckBox::stateChanged, this, &ParserWidget::textEdited); - connect(m_ui->languageStandardsC, static_cast(&QComboBox::activated), this, -&ParserWidget::languageStandardChangedC); - connect(m_ui->languageStandardsCpp, static_cast(&QComboBox::activated), this, -&ParserWidget::languageStandardChangedCpp); + connect(m_ui->languageStandardsC, static_cast(&QComboBox::activated), + this, &ParserWidget::languageStandardChangedC); + connect(m_ui->languageStandardsCpp, static_cast(&QComboBox::activated), + this, &ParserWidget::languageStandardChangedCpp); + connect(m_ui->languageStandardsOpenCl, static_cast(&QComboBox::activated), + this, &ParserWidget::languageStandardChangedOpenCl); + connect(m_ui->languageStandardsCuda, static_cast(&QComboBox::activated), + this, &ParserWidget::languageStandardChangedCuda); updateEnablements(); } @@ -117,6 +121,34 @@ updateEnablements(); } +void ParserWidget::languageStandardChangedOpenCl(const QString& standard) +{ + if (m_ui->languageStandardsOpenCl->currentIndex() == customProfileIdx) { + m_ui->parserOptionsOpenCl->setText(SettingsManager::globalInstance()->defaultParserArguments().openClArguments); + } else { + auto text = SettingsManager::globalInstance()->defaultParserArguments().openClArguments; + auto currentStandard = languageStandard(text); + m_ui->parserOptionsOpenCl->setText(text.replace(currentStandard, standard)); + } + + textEdited(); + updateEnablements(); +} + +void ParserWidget::languageStandardChangedCuda(const QString& standard) +{ + if (m_ui->languageStandardsCuda->currentIndex() == customProfileIdx) { + m_ui->parserOptionsCuda->setText(SettingsManager::globalInstance()->defaultParserArguments().cudaArguments); + } else { + auto text = SettingsManager::globalInstance()->defaultParserArguments().cudaArguments; + auto currentStandard = languageStandard(text); + m_ui->parserOptionsCuda->setText(text.replace(currentStandard, standard)); + } + + textEdited(); + updateEnablements(); +} + void ParserWidget::setParserArguments(const ParserArguments& arguments) { auto setArguments = [this](QComboBox* languageStandards, QLineEdit* parserOptions, const QString& arguments) { @@ -136,28 +168,29 @@ setArguments(m_ui->languageStandardsCpp, m_ui->parserOptionsCpp, arguments.cppArguments); setArguments(m_ui->languageStandardsC, m_ui->parserOptionsC, arguments.cArguments); + setArguments(m_ui->languageStandardsOpenCl, m_ui->parserOptionsOpenCl, arguments.openClArguments); + setArguments(m_ui->languageStandardsCuda, m_ui->parserOptionsCuda, arguments.cudaArguments); m_ui->parseHeadersInPlainC->setChecked(!arguments.parseAmbiguousAsCPP); updateEnablements(); } ParserArguments ParserWidget::parserArguments() const { - return {m_ui->parserOptionsC->text(), m_ui->parserOptionsCpp->text(), {}, !m_ui->parseHeadersInPlainC->isChecked()}; + return { + m_ui->parserOptionsC->text(), + m_ui->parserOptionsCpp->text(), + m_ui->parserOptionsOpenCl->text(), + m_ui->parserOptionsCuda->text(), + !m_ui->parseHeadersInPlainC->isChecked() + }; } void ParserWidget::updateEnablements() { - if (m_ui->languageStandardsCpp->currentIndex() == customProfileIdx) { - m_ui->parserOptionsCpp->setEnabled(true); - } else { - m_ui->parserOptionsCpp->setEnabled(false); - } - - if (m_ui->languageStandardsC->currentIndex() == customProfileIdx) { - m_ui->parserOptionsC->setEnabled(true); - } else { - m_ui->parserOptionsC->setEnabled(false); - } + m_ui->parserOptionsCpp->setEnabled(m_ui->languageStandardsCpp->currentIndex() == customProfileIdx); + m_ui->parserOptionsC->setEnabled(m_ui->languageStandardsC->currentIndex() == customProfileIdx); + m_ui->parserOptionsOpenCl->setEnabled(m_ui->languageStandardsOpenCl->currentIndex() == customProfileIdx); + m_ui->parserOptionsCuda->setEnabled(m_ui->languageStandardsCuda->currentIndex() == customProfileIdx); } diff --git a/languages/plugins/custom-definesandincludes/kcm_widget/parserwidget.ui b/languages/plugins/custom-definesandincludes/kcm_widget/parserwidget.ui --- a/languages/plugins/custom-definesandincludes/kcm_widget/parserwidget.ui +++ b/languages/plugins/custom-definesandincludes/kcm_widget/parserwidget.ui @@ -104,13 +104,26 @@ + + + + Qt::Vertical + + + + 20 + 40 + + + + C Profi&le: - languageStandardsCpp + languageStandardsC @@ -183,8 +196,8 @@ - - + + Qt::Vertical @@ -196,6 +209,167 @@ + + + + OpenCL &Profile: + + + languageStandardsOpenCl + + + + + + + + 0 + 0 + + + + + 100 + 0 + + + + <html><head/><body><p>Choose language profile. </p><p>Use &quot;Custom&quot; profile to modify parser command-line arguments</p></body></html> + + + + Custom + + + + + c++03 + + + + + c++11 + + + + + c++14 + + + + + c99 + + + + + + + + OpenCL comma&nd-line arguments: + + + parserOptionsOpenCl + + + + + + + + 100 + 0 + + + + + + + + Qt::Vertical + + + + 20 + 40 + + + + + + + + CUDA C P&rofile: + + + languageStandardsCuda + + + + + + + + 0 + 0 + + + + + 100 + 0 + + + + <html><head/><body><p>Choose language profile. </p><p>Use &quot;Custom&quot; profile to modify parser command-line arguments</p></body></html> + + + + Custom + + + + + c++03 + + + + + c++11 + + + + + c++14 + + + + + c99 + + + + + + + + CUDA C comman&d-line arguments: + + + parserOptionsCuda + + + + + + + + 100 + 0 + + + +