- c - 在位数组中找到第一个零
- linux - Unix 显示有关匹配两种模式之一的文件的信息
- 正则表达式替换多个文件
- linux - 隐藏来自 xtrace 的命令
我正在开发一个应用程序,它将一个字符串分成多个部分并将每个分配给一个 block 。在每个 block 中,文本逐个字符和一个共享的 int 数组进行扫描,D 将根据读取的字符由不同的线程并行更新。在每次迭代结束时检查 D 的最后一个元素,如果它满足条件,则在与文本对应的位置将全局 int 数组 m 设置为 1。此代码在 NVIDIA GEForce Fermi 550 上执行,运行速度比 CPU 版本还要慢。我刚刚在这里包含了内核:
__global__ void match(uint32_t* BB_d,const char* text_d,int n, int m,int k,int J,int lc,int start_addr,int tBlockSize,int overlap ,int* matched){
__shared__ int D[MAX_THREADS+2];
__shared__ char Text_S[MAX_PATTERN_SIZE];
__shared__ int DNew[MAX_THREADS+2];
__shared__ int BB_S[4][MAX_THREADS];
int w=threadIdx.x+1;
for(int i=0;i<4;i++)
{
BB_S[i][threadIdx.x]= BB_d[i*J+threadIdx.x];
}
{
D[threadIdx.x] = 0;
{
D[w] = (1<<(k+1)) -1;
for(int i = 0; i < lc - 1; i++)
{
D[w] = (D[w] << k+2) + (1<<(k+1)) -1;
}
}
D[J+1] = (1<<((k+2)*lc)) - 1;
}
int startblock=(blockIdx.x == 0?start_addr:(start_addr+(blockIdx.x * (tBlockSize-overlap))));
int size= (((startblock + tBlockSize) > n )? ((n- (startblock))):( tBlockSize));
int copyBlock=(size/J)+ ((size%J)==0?0:1);
if((threadIdx.x * copyBlock) <= size)
memcpy(Text_S+(threadIdx.x*copyBlock),text_d+(startblock+threadIdx.x*copyBlock),(((((threadIdx.x*copyBlock))+copyBlock) > size)?(size-(threadIdx.x*copyBlock)):copyBlock));
memcpy(DNew, D, (J+2)*sizeof(int));
__syncthreads();
uint32_t initial = D[1];
uint32_t x;
uint32_t mask = 1;
for(int i = 0; i < lc - 1; i++)mask = (mask<<(k+2)) + 1;
for(int i = 0; i < size;i++)
{
{
x = ((D[w] >> (k+2)) | (D[w - 1] << ((k + 2)* (lc - 1))) | (BB_S[(((int)Text_S[i])/2)%4][w-1])) & ((1 << (k + 2)* lc) - 1);
DNew[w] = ((D[w]<<1) | mask)
& (((D[w] << k+3) | mask|((D[w +1] >>((k+2)*(lc - 1)))<<1)))
& (((x + mask) ^ x) >> 1)
& initial;
}
__syncthreads();
memcpy(D, DNew, (J+2)*sizeof(int));
if(!(D[J] & 1<<(k + (k + 2)*(lc*J -m + k ))))
{
matched[startblock+i] = 1;
D[J] |= ((1<<(k + 1 + (k + 2)*(lc*J -m + k ))) - 1);
}
}
}
我对CUDA不是很熟悉,所以不太了解共享内存库冲突等问题。这可能是这里的瓶颈吗?
根据要求,这是我启动内核的代码:
#include <stdio.h>
#include <assert.h>
#include <cuda.h>
#define uint32_t unsigned int
#define MAX_THREADS 512
#define MAX_PATTERN_SIZE 1024
#define MAX_BLOCKS 8
#define MAX_STREAMS 16
#define TEXT_MAX_LENGTH 1000000000
void calculateBBArray(uint32_t** BB,const char* pattern_h,int m,int k , int lc , int J){};
void checkCUDAError(const char *msg) {
cudaError_t err = cudaGetLastError();
if( cudaSuccess != err)
{
fprintf(stderr, "Cuda error: %s: %s.\n", msg,
cudaGetErrorString( err) );
exit(EXIT_FAILURE);
}
}
char* getTextString() {
FILE *input, *output;
char c;
char * inputbuffer=(char *)malloc(sizeof(char)*TEXT_MAX_LENGTH);
int numchars = 0, index = 0;
input = fopen("sequence.fasta", "r");
c = fgetc(input);
while(c != EOF)
{
inputbuffer[numchars] = c;
numchars++;
c = fgetc(input);
}
fclose(input);
inputbuffer[numchars] = '\0';
return inputbuffer;
}
int main(void) {
const char pattern_h[] = "TACACGAGGAGAGGAGAAGAACAACGCGACAGCAGCAGACTTTTTTTTTTTTACAC";
char * text_h=getTextString(); //reading text from file, supported upto 200MB currently
int k = 13;
int i;
int count=0;
char *pattern_d, *text_d; // pointers to device memory
char* text_new_d;
int* matched_d;
int* matched_new_d;
uint32_t* BB_d;
uint32_t* BB_new_d;
int* matched_h = (int*)malloc(sizeof(int)* strlen(text_h));
cudaMalloc((void **) &pattern_d, sizeof(char)*strlen(pattern_h)+1);
cudaMalloc((void **) &text_d, sizeof(char)*strlen(text_h)+1);
cudaMalloc((void **) &matched_d, sizeof(int)*strlen(text_h));
cudaMemcpy(pattern_d, pattern_h, sizeof(char)*strlen(pattern_h)+1, cudaMemcpyHostToDevice);
cudaMemcpy(text_d, text_h, sizeof(char)*strlen(text_h)+1, cudaMemcpyHostToDevice);
cudaMemset(matched_d, 0,sizeof(int)*strlen(text_h));
int m = strlen(pattern_h);
int n = strlen(text_h);
uint32_t* BB_h[4];
unsigned int maxLc = ((((m-k)*(k+2)) > (31))?(31/(k+2)):(m-k));
unsigned int lc=2; // Determines the number of threads per block
// can be varied upto maxLc for tuning performance
if(lc>maxLc)
{
exit(0);
}
unsigned int noWordorNfa =((m-k)/lc) + (((m-k)%lc) == 0?0:1);
cudaMalloc((void **) &BB_d, sizeof(int)*noWordorNfa*4);
if(noWordorNfa >= MAX_THREADS)
{
printf("Error: max threads\n");
exit(0);
}
calculateBBArray(BB_h,pattern_h,m,k,lc,noWordorNfa); // not included this function
for(i=0;i<4;i++)
{
cudaMemcpy(BB_d+ i*noWordorNfa, BB_h[i], sizeof(int)*noWordorNfa, cudaMemcpyHostToDevice);
}
int overlap=m;
int textBlockSize=(((m+k+1)>n)?n:(m+k+1));
cudaStream_t stream[MAX_STREAMS];
for(i=0;i<MAX_STREAMS;i++) {
cudaStreamCreate( &stream[i] );
}
int start_addr=0,index=0,maxNoBlocks=0;
if(textBlockSize>n)
{
maxNoBlocks=1;
}
else
{
maxNoBlocks=((1 + ((n-textBlockSize)/(textBlockSize-overlap)) + (((n-textBlockSize)%(textBlockSize-overlap)) == 0?0:1)));
}
int kernelBlocks = ((maxNoBlocks > MAX_BLOCKS)?MAX_BLOCKS:maxNoBlocks);
int blocksRemaining =maxNoBlocks;
printf(" maxNoBlocks %d kernel Blocks %d \n",maxNoBlocks,kernelBlocks);
while(blocksRemaining >0)
{
kernelBlocks = ((blocksRemaining > MAX_BLOCKS)?MAX_BLOCKS:blocksRemaining);
printf(" Calling %d Blocks with starting Address %d , textBlockSize %d \n",kernelBlocks,start_addr,textBlockSize);
match<<<kernelBlocks,noWordorNfa,0,stream[(index++)%MAX_STREAMS]>>>(BB_d,text_d,n,m,k,noWordorNfa,lc,start_addr,textBlockSize,overlap,matched_d);
start_addr+=kernelBlocks*(textBlockSize-overlap);;
blocksRemaining -= kernelBlocks;
}
cudaMemcpy(matched_h, matched_d, sizeof(int)*strlen(text_h), cudaMemcpyDeviceToHost);
checkCUDAError("Matched Function");
for(i=0;i<MAX_STREAMS;i++)
cudaStreamSynchronize( stream[i] );
// do stuff with matched
// ....
// ....
free(matched_h);cudaFree(pattern_d);cudaFree(text_d);cudaFree(matched_d);
return 0;
每个 block 启动的线程数取决于 pattern_h 的长度(最多可以是上面的 maxLc)。在这种情况下,我预计它会在 30 左右。难道这还不足以看到大量的并发吗?至于 block ,我认为一次启动超过 MAX_BLOCKS (=10) 个是没有意义的,因为硬件只能同时调度 8 个
注意:我没有 GUI 访问权限。
最佳答案
对于你正在使用的所有共享内存,如果连续的线程没有从共享数组中的连续地址读取,你可能会遇到存储体冲突......这可能导致内存访问序列化,这反过来会杀死算法的并行性能。
关于c - GPU 代码运行速度低于 CPU 版本,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/13592194/
我想知道在谈到 CPU 使用率和 CPU 利用率时,术语是否存在科学差异。我觉得这两个词都被用作同义词。它们都描述了 CPU 时间和 CPU 容量之间的关系。 Wikipedia称之为 CPU 使用率
我研究了一些关于处理器和 Tomasulo 算法的指令重新排序的内容。 为了更深入地了解这个主题,我想知道是否有任何方法可以(获取跟踪)查看为给定程序完成的实际动态重新排序? 我想给出一个输入程序并查
我有一台配备 2 个 Intel Xeon CPU E5-2620 (Sandy Bridge) 和 10Gbps 82599 NIC(2 个端口)的服务器,用于高性能计算。从 PCI 关联性中,我看
您能详细解释一下“用户 CPU 时间”和“系统 CPU 时间”吗?我读了很多,但我不太理解。 最佳答案 区别在于时间花在用户空间还是内核空间。用户 CPU 时间是处理器运行程序代码(或库中的代码)所花
我想知道如何识别 CPU 是否与 ARM v5 指令集兼容。 假设 ARM v7 指令与 ARM v5 兼容是否正确? 最佳答案 您可以阅读 CPUID base register获得PARTNO。然
我目前在具有多个六核 CPU 的服务器上使用 C 多线程。我想将我的一些线程的亲和性设置为单个 CPU 的各个核心。我使用过 pthread_setaffinity_np() 和 sched_seta
1) 独占时间是在方法中花费的时间2) 包含时间是在方法中花费的时间加上在任何被调用函数中花费的时间3)我们称调用方法为“ parent ”,称方法为“子”。引用链接:Click here 这里的问题
关闭。这个问题需要多问focused 。目前不接受答案。 想要改进此问题吗?更新问题,使其仅关注一个问题 editing this post . 已关闭 5 年前。 Improve this ques
好的,所以编译器可以出于性能原因自由地重新排序代码片段。让我们假设一些代码片段,在没有应用优化的情况下直接翻译成机器代码,看起来像这样: machine_instruction_1 machine_i
我在 zabbix 中有以下默认图表,但我不知道如何解释这些值。谁能解释一下? 最佳答案 操作系统是一件非常忙碌的事情,尤其是当你让它做某事时(即使你没有做)。当我们看到一个活跃的企业环境时,总会发生
换句话说,L1、L2、L3 等缓存是否总是反射(reflect) CPU的字节序 ? 或者总是将数据存储在某些 的缓存中更有意义吗?特定字节序 ? 有没有总体设计决策 ? 最佳答案 大多数现代缓存不会
我想知道当前的 cpus 是否避免在其中至少一个为零时将两个数字相乘。谢谢 最佳答案 这取决于 CPU 和(在某些情况下)操作数的类型。 较旧/较简单的 CPU 通常使用如下乘法算法: integer
我有一个 CUDA 应用程序,它在一台计算机(配备 GTX 275)上运行良好,而在另一台配备 GeForce 8400 的计算机上运行速度慢了大约 100 倍。我怀疑有某种回退使代码实际上在 CPU
例如,对于 8 位 CPU,堆栈大小预计为 8 位宽,16 位 CPU 与 16 位堆栈宽度,以及 32 位、64 位 CPU,等等。是否适用于所有架构? 最佳答案 CPU 具有数据总线和地址总线。它
实现 SIMD 是否需要多核 CPU? 在阅读有关 SIMD 的维基百科时,我发现了以下短语“多处理元素”。那么这句话和“多核CPU”有什么区别呢? 最佳答案 不,每个内核通常都可以执行指令集中的大多
我遗漏了一些基本的东西。 CPU 流水线:在基本层面上,为什么指令需要不同数量的时钟周期才能完成,为什么有些指令在多级 CPU 中只需要 1 个周期? 除了明显的“不同的指令需要不同的工作量才能完成”
超线程 CPU 是实现并行还是仅实现并发(上下文切换)? 我的猜测是没有并行性,只有通过上下文切换的并发性。 最佳答案 单个物理 CPU 具有超线程的核心显示为 两个逻辑 CPU 到操作系统。 CPU
关闭。这个问题不符合Stack Overflow guidelines .它目前不接受答案。 这个问题似乎不是关于 a specific programming problem, a softwar
背景是这样的:下周我们的办公室将有一天因为维护而没有暖气。预计室外温度在 7 至 12 摄氏度之间,因此可能会变冷。可移植电取暖器数量太少,无法满足所有人的需求。 但是,在我大约 6-8 平方米的办公
我开发了一个应用程序,该应用程序在我的开发箱上的三个容器中运行,该开发箱具有带超线程的四核,这意味着系统和 docker 使用 8 个核心。 容器的 CPU 分配由 docker-compose 完成
我是一名优秀的程序员,十分优秀!